/*
 *  Copyright 2009 Advanced Industrial Science and Technology (AIST).
 *
 *  Licensed under the Apache License, Version 2.0 (the "License");
 *  you may not use this file except in compliance with the License.
 *  You may obtain a copy of the License at
 *
 *      http://www.apache.org/licenses/LICENSE-2.0
 *
 *  Unless required by applicable law or agreed to in writing, software
 *  distributed under the License is distributed on an "AS IS" BASIS,
 *  WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
 *  See the License for the specific language governing permissions and
 *  limitations under the License.
 */

/*! \file
 *  \brief
 */
#pragma once

// thrust, includes
#include <thrust/mpl/if.hpp>
#include <thrust/pending/integer_range.hpp>
#include <thrust/pending/functional.hpp>
#include <thrust/graph/detail/adjacency_iterator.hpp>
#include <thrust/graph/detail/edge_iterator.hpp>
#include <thrust/graph/detail/undirected_out_edge_iterator.hpp>
#include <thrust/graph/detail/vertex.hpp>

#include <thrust/device_vector.h>
#include <thrust/device_new.h>
#include <thrust/device_delete.h>
#include <thrust/pair.h>
#include <thrust/tuple.h>
#include <thrust/iterator/zip_iterator.h>
#include <thrust/sort.h>
#include <thrust/count.h>
#include <thrust/binary_search.h>

// stl, includes
#include <iostream>

// system, includes
#include <cassert>

namespace thrust {
  
  namespace detail {
    
    template <typename DirectedS>
    struct directed_category_traits {
      typedef directed_tag directed_category;
    };
    
    template <>
    struct directed_category_traits<directedS> {
      typedef directed_tag directed_category;
    };
    
    template <>
    struct directed_category_traits<undirectedS> {
      typedef undirected_tag directed_category;
    };
    
    template <>
    struct directed_category_traits<bidirectionalS> {
      typedef bidirectional_tag directed_category;
    };
    
    template <typename ParallelS>
    struct edge_parallel_category_traits {
      typedef allow_parallel_edge_tag edge_parallel_category;
    };
    
    template <>
    struct edge_parallel_category_traits<allow_parallelS> {
      typedef allow_parallel_edge_tag edge_parallel_category;
    };
    
    template <>
    struct edge_parallel_category_traits<disallow_parallelS> {
      typedef disallow_parallel_edge_tag edge_parallel_category;
    };
  } // end namespace detail
  
  //----------------------------
  // Directed Edges Helper Class
  
  template <typename Config>
  struct directed_edges_helper {};

  /*!
   *
   */
  template <typename Config>
  __host__
  inline typename Config::degree_size_type
  out_degree( typename Config::vertex_descriptor u,
              const directed_edges_helper<Config>& g_);
  
  /*!
   *
   */
  template <typename Config>
  __host__
  inline typename Config::degree_size_type
  out_degree( typename Config::vertex_descriptor u,
              const directed_edges_helper<Config>& g_,
              thrust::called_from_host_tag);
  
  /*!
   *
   */
  template <typename Config>
  __device__
  inline typename Config::degree_size_type
  out_degree( typename Config::vertex_descriptor u,
              const directed_edges_helper<Config>& g_,
              thrust::called_from_device_tag);
  
  /*!
   *
   */
  template <typename Config>
  __host__
  inline thrust::pair<typename Config::out_edge_iterator,
                      typename Config::out_edge_iterator>
  out_edges( typename Config::vertex_descriptor u,
             directed_edges_helper<Config>& g_);

  /*!
   *
   */
  template <typename Config>
  __host__
  inline thrust::pair<typename Config::out_edge_iterator,
                      typename Config::out_edge_iterator>
  out_edges( typename Config::vertex_descriptor u,
             directed_edges_helper<Config>& g_,
             thrust::called_from_host_tag tag);

  /*!
   *
   */
  template <typename Config>
  __device__
  inline thrust::pair<typename Config::kernel_out_edge_iterator,
                      typename Config::kernel_out_edge_iterator>
  out_edges( typename Config::vertex_descriptor u,
             directed_edges_helper<Config>& g_,
             thrust::called_from_device_tag tag);
  
  //=============================
  // Directed Graph Helper Class
  
  struct adj_list_dir_traversal_tag :
    public virtual vertex_list_graph_tag
    , public virtual incidence_graph_tag
    , public virtual adjacency_graph_tag
    , public virtual edge_list_graph_tag {};
  
  template <typename Config>
  struct directed_graph_helper 
    : public directed_edges_helper<Config>
  {
    
    typedef adj_list_dir_traversal_tag traversal_category;
    
    __host__ 
    inline void 
    init( typename Config::edges_size_type capacity_edges)
    {
      typedef typename Config::graph_type         Graph;
      typedef typename Config::edge_descriptor    edge_descriptor;
      typedef typename Config::edge_property_type edge_property_type;
      typedef typename Config::EdgeContainer         EdgeContainer;
      typedef typename Config::EdgePropertyContainer EdgePropertyContainer;
      
      Graph* g =static_cast<Graph*>(this);
      
      g->m_capacity_out_edges =capacity_edges;
			
      g->m_out_edges.resize( capacity_edges);
      g->ptr_out_edges =thrust::device_pointer_cast<edge_descriptor>( &g->m_out_edges[0]);
      
      if( !thrust::detail::same_property<edge_property_type,no_property>::value)
        g->m_edge_property.resize( capacity_edges);
    }
  };
	
  /*!
   *
   */
  template <typename Config>
  __host__ 
  inline thrust::pair<typename Config::edge_descriptor, bool>
  add_edge( typename Config::vertex_descriptor u,
            typename Config::vertex_descriptor v,
            directed_graph_helper<Config>& g_);
  
  /*!
   *
   */
  template <typename Config>
  __host__
  inline void
  remove_edge( typename Config::edge_descriptor e ,
               directed_graph_helper<Config>& g_,
               bool erase_edge_flag =false);
  
  //===============================
  // UnDirected Graph Helper Class
  
  struct undir_adj_list_traversal_tag :
    public virtual vertex_list_graph_tag
    , public virtual incidence_graph_tag
    , public virtual adjacency_graph_tag
    , public virtual edge_list_graph_tag
    , public virtual bidirectional_graph_tag {};
  
  template <typename Config>
  struct undirected_graph_helper
  {
    typedef undir_adj_list_traversal_tag traversal_category;
		
    __host__ inline
    void init( typename Config::edges_size_type capacity_edges)
    {
      typedef typename Config::graph_type Graph;
      typedef typename Config::edge_descriptor    edge_descriptor;
      typedef typename Config::edge_property_type edge_property_type;
      typedef typename Config::EdgeContainer         EdgeContainer;
      typedef typename Config::EdgePropertyContainer EdgePropertyContainer;
			
      Graph* g =static_cast<Graph*>(this);
			
      g->m_capacity_out_edges =capacity_edges;
      g->m_capacity_in_edges  =capacity_edges;
      
      g->m_out_edges.resize( capacity_edges);
      g->m_in_edges.resize( capacity_edges);
      
      g->ptr_out_edges =thrust::device_pointer_cast<edge_descriptor>( &g->m_out_edges[0]);
      g->ptr_in_edges  =thrust::device_pointer_cast<edge_descriptor>( &g->m_in_edges[0]);
      
      if( !thrust::detail::same_property<edge_property_type,no_property>::value)
        g->m_edge_property.resize( capacity_edges);
    }
  };
	
  namespace detail
  {
    
    template <typename Config>
    inline void
    reconstruct_in_edges( undirected_graph_helper<Config>& g_
                          , typename Config::vertices_size_type V
                          , typename Config::edges_size_type    E);

    /*!
     *
     */
    template <typename Config>
    __host__
    inline typename Config::degree_size_type
    out_degree( typename Config::vertex_descriptor u,
                const undirected_graph_helper<Config>& g_);
    
    /*!
     *
     */
    template <typename Config>
    __host__
    inline typename Config::degree_size_type
    in_degree( typename Config::vertex_descriptor u,
               const undirected_graph_helper<Config>& g_);
 
    /*!
     *
     */
    template <typename Config>
    __host__
    inline thrust::pair<typename Config::edge_iterator,
                        typename Config::edge_iterator>
    out_edges( typename Config::vertex_descriptor u,
               undirected_graph_helper<Config>& g_);
		
    /*!
     *
     */
    template <typename Config>
    __host__
    inline thrust::pair<typename Config::edge_iterator,
												typename Config::edge_iterator>
    in_edges( typename Config::vertex_descriptor u,
              undirected_graph_helper<Config>& g_);
		
    /*!
     *
     */
    template <typename Config>
    __host__
    inline thrust::pair<typename Config::edge_iterator,
                        typename Config::edge_iterator>
    in_edges( undirected_graph_helper<Config>& g_);
    
  } // end namespace detail
	
  /*!
   *
   */
  template <typename Config>
  __host__
  inline typename Config::degree_size_type
  out_degree( typename Config::vertex_descriptor u,
              const undirected_graph_helper<Config>& g_);
  
  /*!
   *
   */    
  template <typename Config>
  __host__
  inline typename Config::degree_size_type
  out_degree( typename Config::vertex_descriptor u,
              const undirected_graph_helper<Config>& g_,
              thrust::called_from_host_tag);
	
  /*!
   *
   */
  template <typename Config>
  __device__
  inline typename Config::degree_size_type
  out_degree( typename Config::vertex_descriptor u,
              const undirected_graph_helper<Config>& g_,
              thrust::called_from_device_tag);
  
  /*!
   *
   */
  template <typename Config>
  __host__
  inline thrust::pair<typename Config::out_edge_iterator,
                      typename Config::out_edge_iterator>
  out_edges( typename Config::vertex_descriptor u,
             undirected_graph_helper<Config>& g_);
  
  /*!
   *
   */
  template <typename Config>
  __host__
  inline thrust::pair<typename Config::out_edge_iterator,
                      typename Config::out_edge_iterator>
  out_edges( typename Config::vertex_descriptor u,
             undirected_graph_helper<Config>& g_,
             thrust::called_from_host_tag tag);

  /*!
   *
   */
  template <typename Config>
  __device__
  inline thrust::pair<typename Config::kernel_out_edge_iterator,
                      typename Config::kernel_out_edge_iterator>
  out_edges( typename Config::vertex_descriptor u,
             undirected_graph_helper<Config>& g_,
             thrust::called_from_device_tag tag);

  /*!
   *
   */
  template <typename Config>
  __host__ 
  inline thrust::pair<typename Config::edge_descriptor, bool>
  add_edge( typename Config::vertex_descriptor u,
            typename Config::vertex_descriptor v,
            undirected_graph_helper<Config>& g_);

  /*!
   *
   */
  template <typename Config>
  __host__
  inline void
  remove_edge( typename Config::edge_descriptor e,
               undirected_graph_helper<Config>& g_,
               bool erase_edge_flag =false);
  
  //==================================
  // Bidirectional Graph Helper Class
  
  struct bidir_adj_list_traversal_tag :
    public virtual vertex_list_graph_tag
    , public virtual incidence_graph_tag
    , public virtual adjacency_graph_tag
    , public virtual edge_list_graph_tag
    , public virtual bidirectional_graph_tag {};
  
  template <typename Config>
  struct bidirectional_graph_helper 
    : public directed_edges_helper<Config>
  {
    typedef bidir_adj_list_traversal_tag traversal_category;
		
    __host__ inline
    void init( typename Config::edges_size_type capacity_edges){
			
      typedef typename Config::graph_type         Graph;
      typedef typename Config::edge_descriptor    edge_descriptor;
      typedef typename Config::edge_property_type edge_property_type;
      typedef typename Config::EdgeContainer         EdgeContainer;
      typedef typename Config::EdgePropertyContainer EdgePropertyContainer;
      
      Graph* g =static_cast<Graph*>(this);
			
      g->m_capacity_out_edges =capacity_edges;
      g->m_capacity_in_edges  =capacity_edges;
			
      g->m_out_edges.resize( capacity_edges);
      g->m_in_edges.resize( capacity_edges);
			
      g->ptr_out_edges =thrust::device_pointer_cast<edge_descriptor>( &g->m_out_edges[0]);
      g->ptr_in_edges  =thrust::device_pointer_cast<edge_descriptor>( &g->m_in_edges[0]);
			
      if( !thrust::detail::same_property<edge_property_type,no_property>::value)
        g->m_edge_property.resize( capacity_edges);      
    }
  };
	
  namespace detail {
    template <typename Config>
    inline void
    reconstruct_in_edges( bidirectional_graph_helper<Config>& g_
                          , typename Config::vertices_size_type V
                          , typename Config::edges_size_type    E);
		
    /*!
     *
     */
    template <typename Config>
    __host__
    inline thrust::pair<typename Config::edge_iterator,
                        typename Config::edge_iterator>
    in_edges( bidirectional_graph_helper<Config>& g_);

  } // end namespace detail
	
  /*!
   *
   */
  template <typename Config>
  __host__
  inline typename Config::degree_size_type
  in_degree( typename Config::vertex_descriptor u,
             const bidirectional_graph_helper<Config>& g_);
	
  /*!
   *
   */
  template <typename Config>
  __host__
  inline typename Config::degree_size_type
  in_degree( typename Config::vertex_descriptor u,
             const bidirectional_graph_helper<Config>& g_,
             thrust::called_from_host_tag);
	
  /*!
   *
   */
  template <typename Config>
  __device__
  inline typename Config::degree_size_type
  in_degree( typename Config::vertex_descriptor u,
             const bidirectional_graph_helper<Config>& g_,
             thrust::called_from_device_tag);
	
  /*!
   *
   */
  template <typename Config>
  __host__
  inline thrust::pair<typename Config::in_edge_iterator,
                      typename Config::in_edge_iterator>
  in_edges( typename Config::vertex_descriptor u,
            bidirectional_graph_helper<Config>& g_);
	
  /*!
   *
   */
  template <typename Config>
  __host__
  inline thrust::pair<typename Config::in_edge_iterator,
                      typename Config::in_edge_iterator>
  in_edges( typename Config::vertex_descriptor u,
            bidirectional_graph_helper<Config>& g_,
            thrust::called_from_host_tag tag);
	
  /*!
   *
   */
  template <typename Config>
  __device__
  inline thrust::pair<typename Config::kernel_in_edge_iterator,
                      typename Config::kernel_in_edge_iterator>
  in_edges( typename Config::vertex_descriptor u,
            bidirectional_graph_helper<Config>& g_,
            thrust::called_from_device_tag tag);
	
  /*!
   *
   */
  template <typename Config>
  __host__
  inline typename Config::degree_size_type
  degree( typename Config::vertex_descriptor u,
          const bidirectional_graph_helper<Config>& g_);
	
  /*!
   *
   */
  template <typename Config>
  __host__
  inline typename Config::degree_size_type
  degree( typename Config::vertex_descriptor u,
          const bidirectional_graph_helper<Config>& g_,
          thrust::called_from_host_tag);
	
  /*!
   *
   */
  template <typename Config>
  __device__
  inline typename Config::degree_size_type
  degree( typename Config::vertex_descriptor u,
          const bidirectional_graph_helper<Config>& g_,
          thrust::called_from_device_tag);
  
  /*!
   *
   */
  template <typename Config>
  __host__ 
  inline thrust::pair<typename Config::edge_descriptor, bool>
  add_edge( typename Config::vertex_descriptor u,
            typename Config::vertex_descriptor v,
            bidirectional_graph_helper<Config>& g_);
	
  /*!
   *
   */
  template <typename Config>
  __host__
  inline void
  remove_edge( typename Config::edge_descriptor e,
               bidirectional_graph_helper<Config>& g_,
               bool erase_edge_flag =false);

  //=============================
  // Adjacency List Helper Class
  template <typename Config, typename DirectedHelper>
  struct adj_list_helper : public DirectedHelper
  {
    __host__ inline
    void init( typename Config::edges_size_type capacity_edges)
    {
      DirectedHelper* g =static_cast<DirectedHelper*>(this);
      g->init( capacity_edges);
    }
  };
  
  namespace detail {
    template <typename Config, typename DirectedHelper>
    inline void
    reconstruct_in_edges( adj_list_helper<Config,DirectedHelper>& g_
                          , typename Config::vertices_size_type V
                          , typename Config::edges_size_type    E);
  
    /*!
     *
     */
    template <typename Config, typename DirectedHelper>
    __host__ __device__
    inline typename Config::edges_size_type
    capacity_in_edges( const adj_list_helper<Config,DirectedHelper>& g_);

    /*!
     *
     */
    template <typename Config, typename DirectedHelper>
    __host__ __device__
    inline typename Config::edges_size_type
    num_in_edges( const adj_list_helper<Config,DirectedHelper>& g_);

    /*!
     *
     */
    template <typename Config, typename DirectedHelper>
    __host__
    inline thrust::pair<typename Config::edge_iterator,
                        typename Config::edge_iterator>
    all_out_edges( adj_list_helper<Config,DirectedHelper>& g_);
		
    /*!
     *
     */
    template <typename Config, typename DirectedHelper>
    __host__
    inline thrust::pair<typename Config::edge_iterator,
                        typename Config::edge_iterator>
    all_in_edges( adj_list_helper<Config,DirectedHelper>& g_);
    
  } // end namespace detail

  /*!
   *
   */
  template <typename Config, typename DirectedHelper>
  __host__
  thrust::pair<typename Config::adjacency_iterator,
               typename Config::adjacency_iterator>
  adjacent_vertices( typename Config::vertex_descriptor u,
                     adj_list_helper<Config,DirectedHelper>& g_);

  /*!
   *
   */
  template <typename Config, typename DirectedHelper>
  __device__
  thrust::pair<typename Config::kernel_adjacency_iterator,
               typename Config::kernel_adjacency_iterator>
  adjacent_vertices( typename Config::vertex_descriptor u,
                     adj_list_helper<Config,DirectedHelper>& g_, 
                     thrust::called_from_device_tag tag);
	
  /*!
   *
   */
  template <typename Config, typename DirectedHelper>
  __host__ __device__
  inline typename Config::vertices_size_type
  num_vertices( const adj_list_helper<Config,DirectedHelper>& g_);

  /*!
   *
   */
  template <typename Config, typename DirectedHelper>
  __host__ __device__
  inline typename Config::vertices_size_type
  capacity_vertices( const adj_list_helper<Config,DirectedHelper>& g_);
	
  /*!
   *
   */
  template <typename Config, typename DirectedHelper>
  __host__ __device__
  inline typename Config::edges_size_type
  num_edges( const adj_list_helper<Config,DirectedHelper>& g_);

  /*!
   *
   */
  template <typename Config, typename DirectedHelper>
  __host__ __device__
  inline typename Config::edges_size_type
  capacity_edges( const adj_list_helper<Config,DirectedHelper>& g_);

  /*!
   *
   */
  template <typename Config, typename DirectedHelper>
  __host__ __device__
  inline thrust::pair<typename Config::vertex_iterator,
                      typename Config::vertex_iterator>
  vertices( const adj_list_helper<Config,DirectedHelper>& g_);

  /*!
   *
   */
  template <typename Config, typename DirectedHelper>
  __host__
  inline thrust::pair<typename Config::vertex_iterator,
                      typename Config::vertex_iterator>
  vertices( const adj_list_helper<Config,DirectedHelper>& g_,
            thrust::called_from_host_tag);

  /*!
   *
   */
  template <typename Config, typename DirectedHelper>
  __device__
  inline thrust::pair<typename Config::vertex_iterator,
                      typename Config::vertex_iterator>
  vertices( const adj_list_helper<Config,DirectedHelper>& g_,
            thrust::called_from_device_tag);

  /*!
   *
   */
  template <typename Config, typename DirectedHelper>
  __host__ 
  inline typename Config::vertex_descriptor
  add_vertex( adj_list_helper<Config,DirectedHelper>& g_);

  /*!
   *
   */
  template <typename Config, typename DirectedHelper>
  __host__
  inline void
  clear_vertex( typename Config::vertex_descriptor u,
                adj_list_helper<Config,DirectedHelper>& g_);

  /*!
   *
   */
  template <typename Config, typename DirectedHelper>
  __host__
  inline thrust::pair<typename Config::edge_iterator,
                      typename Config::edge_iterator>
  edges( adj_list_helper<Config,DirectedHelper>& g_);

  /*!
   *
   */
  template <typename Config, typename DirectedHelper>
  __host__
  inline thrust::pair<typename Config::edge_iterator,
                      typename Config::edge_iterator>
  edges( adj_list_helper<Config,DirectedHelper>& g_,
         thrust::called_from_host_tag);

  /*!
   *
   */
  template <typename Config, typename DirectedHelper>
  __device__
  inline thrust::pair<typename Config::kernel_edge_iterator,
                      typename Config::kernel_edge_iterator>
  edges( adj_list_helper<Config,DirectedHelper>& g_,
         thrust::called_from_device_tag);
	
  template <typename Config, typename DirectedHelper>
  __host__
  inline void
  remove_edge( typename Config::vertex_descriptor u,
							 typename Config::vertex_descriptor v,
							 adj_list_helper<Config,DirectedHelper>& g,
							 bool erase_edge_flag =false);
  
  //===============================================
  // Implamantation of the PropertyGraph Interface
	
  /*!
   *
   */
  template <typename Config
						, typename DirectedHelper
	    , typename Property>
  __host__ __device__
  inline typename thrust::property_map<typename Config::graph_type, Property>::type
  get( Property p
       , adj_list_helper<Config,DirectedHelper>& g);
	
  /*!
   *
   */
  template <typename Config
            , typename DirectedHelper
            , typename Property>
  __host__ __device__
  inline typename thrust::property_map<typename Config::graph_type, Property>::const_type
  get( Property p
       , const adj_list_helper<Config,DirectedHelper>& g);
  
  /*!
   *
   */
  template <typename Config
            , typename DirectedHelper
            , typename Property
            , typename Key>
  __host__ __device__
  inline typename thrust::property_traits<
		typename thrust::property_map<typename Config::graph_type, Property>::type>::reference
  get( Property p
       , adj_list_helper<Config,DirectedHelper>& g
       , const Key& key
       , thrust::called_from_device_tag tag);
  
  /*!
   *
   */
  template <typename Config
            , typename DirectedHelper
            , typename Property
            , typename Key>
  __host__ __device__
  inline typename thrust::property_traits<
    typename thrust::property_map<typename Config::graph_type,Property>::const_type>::reference
  get( Property p
       , const adj_list_helper<Config,DirectedHelper>& g
       , const Key& key
       , thrust::called_from_device_tag tag);
  
  /*!
   *
   */
  template <typename Config
            , typename DirectedHelper
            , typename Property
            , typename Key>
  __host__
  inline typename thrust::property_traits<
    typename thrust::property_map<typename Config::graph_type,Property>::type>::reference
  get( Property p
       , adj_list_helper<Config,DirectedHelper>& g
       , const Key& key);
  
  /*!
   *
   */
  template <typename Config
            , typename DirectedHelper
            , typename Property
            , typename Key>
  __host__
  inline typename thrust::property_traits<
    typename thrust::property_map<typename Config::graph_type,Property>::const_type>::reference
  get( Property p
       , const adj_list_helper<Config,DirectedHelper>& g
       , const Key& key);
  
  /*!
   *
   */
  template <typename Config
            , typename DirectedHelper
            , typename Property
            , typename Key
            , typename Value>
  __host__
  inline void
  put( Property p
       , adj_list_helper<Config,DirectedHelper>& g_
       , const Key& key
       , const Value& value
       , thrust::called_from_host_tag);
	
  /*!
   *
   */
  template <typename Config
            , typename DirectedHelper
            , typename Property
            , typename Key
            , typename Value>
  __device__
  inline void
  put( Property p
       , adj_list_helper<Config,DirectedHelper>& g_
       , const Key& key
       , const Value& value
       , thrust::called_from_device_tag);
	
  /*!
   *
   */
  template <typename Config
            , typename DirectedHelper
            , typename Property
            , typename Key
            , typename Value>
  __host__
  inline void
  put( Property p
       , adj_list_helper<Config,DirectedHelper>& g_
       , const Key& key
       , const Value& value);
  
  //======================================
  // Device Adjacency List Implementation
  
  struct adj_list_tag { };
  
  template <typename Graph, typename Config, typename DirectedHelper>
  class adj_list_impl
    : public adj_list_helper<Config,DirectedHelper>
  {
  private:
    typedef adj_list_helper<Config,DirectedHelper> super_t;
    
  public:
    typedef typename Config::vertex_property_type vertex_property_type;
    typedef typename Config::edge_property_type edge_property_type;
    
    /** category tags. **/
    //typedef typename Config::space_tag              space_tag;
    typedef typename Config::directed_category      directed_category;
    typedef typename Config::edge_parallel_category edge_parallel_category;
    
    /** size types. **/
    typedef typename Config::size_type          size_type;
    typedef typename Config::vertices_size_type vertices_size_type;
    typedef typename Config::edges_size_type    edges_size_type;
    typedef typename Config::degree_size_type   degree_size_type;
		
    typedef typename Config::vertex_descriptor vertex_descriptor;
    typedef typename Config::vertex_iterator   vertex_iterator;
    
    typedef typename Config::edge_descriptor   edge_descriptor;
    typedef typename Config::out_edge_iterator out_edge_iterator;
    typedef typename Config::in_edge_iterator  in_edge_iterator;

    typedef typename Config::kernel_out_edge_iterator kernel_out_edge_iterator;
    typedef typename Config::kernel_in_edge_iterator  kernel_in_edge_iterator;

    typedef typename Config::adjacency_iterator        adjacency_iterator;
    typedef typename Config::kernel_adjacency_iterator kernel_adjacency_iterator;

    typedef typename Config::edge_iterator        edge_iterator;
    typedef typename Config::kernel_edge_iterator kernel_edge_iterator;
    
    /** internal value types. **/
    typedef typename Config::VertexContainer VertexContainer; // is device_vector<StoredVertex>
    typedef typename Config::StoredVertex StoredVertex;
    typedef typename Config::EdgeContainer EdgeContainer;     // is device_vector<edge_descriptor>
    typedef typename Config::EdgePropertyContainer EdgePropertyContainer;
  
    typedef adj_list_tag graph_tag;
    
  public:
    vertices_size_type m_num_vertices;
    vertices_size_type m_capacity_vertices;
    edges_size_type    m_num_out_edges;
    edges_size_type    m_num_in_edges;
    edges_size_type    m_capacity_out_edges;
    edges_size_type    m_capacity_in_edges;

    VertexContainer m_vertices;
    EdgeContainer   m_out_edges;
    EdgeContainer   m_in_edges;
    device_vector<edge_property_type> m_edge_property;
    
    device_ptr<StoredVertex> ptr_vertices;
    device_ptr<edge_descriptor> ptr_out_edges;
    device_ptr<edge_descriptor> ptr_in_edges;
		
  public:
    static vertex_descriptor null_vertex()
    {
      return thrust::numeric_limits<vertex_descriptor>::max();
    }
    
  public:
    /*!
     *
     */
    __host__ inline
    adj_list_impl();
    
    /*!
     *
     */
    __host__ inline 
    adj_list_impl( vertices_size_type num_vertices);
    
    /*!
     *
     */
    __host__ inline
    adj_list_impl( vertices_size_type num_vertices,
                   edges_size_type capacity_edges);

    /*!
     *
     */
    __host__ inline
    ~adj_list_impl();
    
  private:
    /*!
     *
     */
    __host__
    inline void
    insert_out_edge( size_type ipos, edge_descriptor& e);
    
    /*!
     *
     */
    __host__
    inline void
    remove_out_edges( edge_iterator lower 
                      , edge_iterator upper
                      , vertex_descriptor u
                      , bool erase_edge_flag);
  public:
    __host__
      inline VertexContainer& 
    get_vertex_container();
    
    __host__
    inline EdgeContainer&
    get_edge_container();
    
    /*!
     *
     */
    __host__ __device__
    inline vertices_size_type
    get_num_vertices() const;
    
    /*!
     *
     */
    __host__ __device__
    inline vertices_size_type
    get_capacity_vertices() const;

    /*!
     *
     */
    __host__ __device__
    inline edges_size_type
    get_num_out_edges() const;

    /*!
     *
     */
    __host__ __device__
    inline edges_size_type
    get_num_in_edges() const;
    
    /*!
     *
     */
    __host__ __device__
    inline edges_size_type
    get_capacity_out_edges() const;

    /*!
     *
     */
    __host__ __device__
    inline edges_size_type
    get_capacity_in_edges() const;

    /*!
     *
     */
    __host__ __device__
    inline integer_range<vertex_descriptor,device_space_tag>
    get_vertices() const;

    /*!
     *
     */
    __host__
    inline  thrust::pair<typename EdgeContainer::iterator,
                         typename EdgeContainer::iterator>
    get_out_edges();
    
    /*!
     *
     */
    __host__
    inline  thrust::pair<typename EdgeContainer::iterator,
                         typename EdgeContainer::iterator>
    get_in_edges();
    
    /*!
     *
     */
    __host__
    inline  thrust::pair<typename EdgeContainer::iterator,
                         typename EdgeContainer::iterator>
    get_all_out_edges();
    
    /*!
     *
     */
    __host__
    inline  thrust::pair<typename EdgeContainer::iterator,
                         typename EdgeContainer::iterator>
    get_all_in_edges();

    //===================================
    // get out and in edges at vertex: u

    /*!
     *
     */
    __host__
    inline thrust::pair<edge_iterator,edge_iterator>
    get_out_edges( vertex_descriptor u
                   , thrust::called_from_host_tag);
    
    /*!
     *
     */
    __device__
    inline thrust::pair<kernel_edge_iterator,kernel_edge_iterator>
    get_out_edges( vertex_descriptor u
                   , thrust::called_from_device_tag);
    
    /*!
     *
     */
    __host__
    inline thrust::pair<edge_iterator,edge_iterator>
    get_in_edges( vertex_descriptor u
                  , thrust::called_from_host_tag);
    
    /*!
     *
     */
    __device__
    inline thrust::pair<kernel_edge_iterator,kernel_edge_iterator>
    get_in_edges( vertex_descriptor u
                  , thrust::called_from_device_tag);
    
    /*!
     *
     */
    __host__
    inline vertex_descriptor
    add_vertex();

    /*!
     *
     */
    __host__
    inline void
    reassign_edge_property();

    /*!
     *
     */
    __host__ 
    inline void
    add_out_edge( edge_descriptor& e);

    /*!
     *
     */
    __host__
    inline void
    remove_out_edge( typename Config::edge_descriptor e,
                     bool erase_edge_flag =false);
  }; // end class adj_list_impl
  
  namespace detail {
    
    template <
      typename Graph
      , typename DirectedS
      , typename ParallelS
      , typename VertexProperty
      , typename EdgeProperty
      , typename SizeType
      >
    struct adj_list_gen {
      
      /*! \p has_property is defined in thrust/pending/property.hpp
       *
       */
      typedef typename has_property<EdgeProperty>::type has_edge_property;
      typedef typename DirectedS::is_directed_t DirectedT;
      typedef typename DirectedS::is_bidir_t    BidirectionalT;
      
      struct config {
        typedef Graph graph_type;
        typedef device_space_tag space_tag;
        
        typedef VertexProperty vertex_property_type;
        typedef EdgeProperty   edge_property_type;
        
        typedef SizeType size_type;
        
        typedef adjacency_list_traits<DirectedS
                                      , ParallelS
                                      , VertexProperty
                                      , EdgeProperty
                                      , SizeType
                                      > Traits;
        
        typedef typename Traits::directed_category      directed_category;
        typedef typename Traits::edge_parallel_category edge_parallel_category;
        
        typedef typename Traits::size_type          vertices_size_type;
        typedef typename Traits::edges_size_type    edges_size_type;
        typedef typename Traits::vertex_descriptor  vertex_descriptor;
        typedef typename Traits::edge_descriptor    edge_descriptor;
        
        typedef vertices_size_type degree_size_type;
        
        //--------------------------------
        // VertexList and vertex_iterator
        typedef integer_range<vertex_descriptor,thrust::device_space_tag> VertexList;
        typedef typename VertexList::iterator vertex_iterator;
        
        typedef typename thrust::mpl::if_<BidirectionalT,
                                          bidir_stored_vertex<size_type,VertexProperty>,
                                          typename thrust::mpl::if_<DirectedT,
                                                                    stored_vertex<size_type,VertexProperty>,
                                                                    bidir_stored_vertex<size_type,VertexProperty>
                                                                    >::type
                                          >::type StoredVertex;
        typedef device_vector<StoredVertex> VertexContainer;
				
        //------------------------------
        // EdgeContainer and edge_iterator
        typedef device_vector<edge_descriptor> EdgeContainer;
        typedef device_vector<edge_property_type> EdgePropertyContainer;
        
        typedef typename EdgeContainer::iterator EdgeContainerIter;
        
        typedef edge_iter<EdgeContainerIter
                          , edge_descriptor
                          , thrust::called_from_host_tag
                          > EdgeIter;
        
        typedef edge_iter<EdgeContainerIter
                          , edge_descriptor
                          , thrust::called_from_device_tag
                          > KernelEdgeIter;
        
        typedef EdgeIter       edge_iterator;
        typedef KernelEdgeIter kernel_edge_iterator;
      
        //------------------------------------------------------------
        // out_edge_iterator, in_edge_iterator and adjacency_iterator
        typedef undirected_out_edge_iter<EdgeContainerIter
                                         , vertex_descriptor
                                         , edge_descriptor
                                         , thrust::called_from_host_tag
                                         > UnDirectedOutEdgeIter;
        
        typedef undirected_out_edge_iter<EdgeContainerIter
                                         , vertex_descriptor
                                         , edge_descriptor
                                         , thrust::called_from_device_tag
                                         > KernelUnDirectedOutEdgeIter;
      
        typedef typename thrust::mpl::if_<DirectedT,
                                          EdgeIter,
                                          UnDirectedOutEdgeIter
                                          >::type out_edge_iterator;
      
        typedef typename thrust::mpl::if_<DirectedT,
                                          KernelEdgeIter,
                                          KernelUnDirectedOutEdgeIter
                                          >::type kernel_out_edge_iterator;
        
        typedef out_edge_iterator in_edge_iterator;
        typedef kernel_out_edge_iterator kernel_in_edge_iterator;
        
        typedef adjacency_iter<graph_type
                               , vertex_descriptor
                               , out_edge_iterator> adjacency_iterator;
        
        typedef adjacency_iter<graph_type
                               , vertex_descriptor
                               , kernel_out_edge_iterator
                               > kernel_adjacency_iterator;

      }; // end config
      
      typedef typename thrust::mpl::if_<BidirectionalT
                                        , bidirectional_graph_helper<config>
                                        , typename thrust::mpl::if_<DirectedT
                                                                    , directed_graph_helper<config>
                                                                    , undirected_graph_helper<config>
                                                                    >::type
                                        >::type DirectedHelper;
      typedef adj_list_impl<Graph,config,DirectedHelper> type;
    }; // end adj_list_gen
  } // end namespace detail
	
  //==========================================================
  // Vertex Property Maps
  // Implemantation of vertex property_map
  template <typename Graph
            , typename ValueType
            , typename Reference
            , typename Tag>
  struct adj_list_vertex_property_map
    : public thrust::put_get_helper< Reference
                                     , adj_list_vertex_property_map<Graph
                                                                    , ValueType
                                                                    , Reference
                                                                    , Tag>
                                     >
  {
  public:
    typedef ValueType value_type;
    typedef Reference reference;
    typedef typename thrust::graph_traits<Graph>::vertex_descriptor key_type;
    typedef typename thrust::lvalue_property_map_tag category;
    typedef typename Graph::StoredVertex StoredVertex;
    
  private:
    Graph* m_g;
    
  public:
    /*!
     *
     */
    __host__ __device__
    adj_list_vertex_property_map();
    
    /*!
     *
     */
    __host__ __device__
    adj_list_vertex_property_map( Graph* g);
    
    /*!
     *
     */
    __host__
    inline Reference
    operator[]( key_type v) const;

    /*!
     *
     */
    __host__
    inline Reference 
    operator()( key_type v) const;
		
    /*!
     *
     */
    __host__ __device__
    inline void
    dispatch_put( key_type k, const value_type& v, thrust::called_from_host_tag);
		
    /*!
     *
     */
    __device__
    inline void
    dispatch_put( key_type k, const value_type& v, thrust::called_from_device_tag);
		
    /*!
     *
     */
    /*template <typename Graph
			, typename ValueType
			, typename Reference
			, typename Tag
			, typename CalledSpace>*/
    template <typename CalledSpace>
    __host__ __device__
    inline void
    put( key_type k, const value_type& v, CalledSpace tag);
  };
  
  // Implemantation of vertex id_map
  template <typename Property, typename Vertex>
  struct adj_list_vertex_id_map
    : public thrust::put_get_helper<Vertex
                                    , adj_list_vertex_id_map<Property, Vertex>
                                    >
  {
    typedef Vertex value_type;
    typedef Vertex key_type;
    typedef Vertex reference;
    typedef thrust::readable_property_map_tag category;
    
    inline adj_list_vertex_id_map() { }
    
    template <typename Graph>
    inline adj_list_vertex_id_map( const Graph&) { }
    
    inline value_type
    operator[]( key_type v) const
    {
      return v;
    }
		
    inline value_type
    operator()( key_type v) const
    {
      return v;
    }
  };
  
  struct adj_list_any_vertex_pa {
    template <typename Tag, typename Graph, typename Property>
    struct bind_ {      
      typedef typename property_value<Property,Tag>::type value_type;
      typedef value_type&       reference;
      typedef const value_type& const_reference;
      
      typedef adj_list_vertex_property_map<Graph
                                           , value_type
                                           , reference
                                           , Tag
                                           > type;
      typedef adj_list_vertex_property_map<Graph
                                           , value_type
                                           , const_reference
                                           , Tag> const_type;
      
    };
  };
  struct adj_list_id_vertex_pa {
    template <typename Tag, typename Graph, typename Property>
    struct bind_ {
      typedef typename Graph::vertex_descriptor Vertex;
      typedef adj_list_vertex_id_map<Property,Vertex> type;
      typedef adj_list_vertex_id_map<Property,Vertex> const_type;
    };
  };
  
  namespace detail {
		
    template <typename Tag>
    struct adj_list_choose_vertex_pa_helper {
      typedef adj_list_any_vertex_pa type;
    };
    template <>
    struct adj_list_choose_vertex_pa_helper<vertex_index_t> {
      typedef adj_list_id_vertex_pa type;
    };
    
    template <typename Tag, typename Graph, typename Property>
    struct adj_list_choose_vertex_pa {
      typedef typename adj_list_choose_vertex_pa_helper<Tag>::type Helper;
      typedef typename Helper::template bind_<Tag,Graph,Property> Bind;
      typedef typename Bind::type       type;
      typedef typename Bind::const_type const_type;
    };
  } //detail
  
  struct adj_list_vertex_property_selector {
    template <typename Graph, typename Property, typename Tag>
    struct bind_ {
      typedef detail::adj_list_choose_vertex_pa<Tag,Graph,Property> Choice;
      typedef typename Choice::type       type;
      typedef typename Choice::const_type const_type;
    };
  };
  
  template <>
  struct vertex_property_selector<adj_list_tag> {
    typedef adj_list_vertex_property_selector type;
  };
  
  //==========================================================
  // Edge Property Maps
  
  // Implemantation of edge property_map
  template <typename Directed, typename ValueType, typename Reference, typename Vertex,
						typename Property, typename Tag>
  struct adj_list_edge_property_map
    : public put_get_helper<Reference
                            , adj_list_edge_property_map<Directed
                                                         , ValueType
                                                         , Reference
                                                         , Vertex
                                                         , Property
                                                         , Tag>
                            >
  {
    typedef ValueType value_type;
    typedef Reference reference;
    typedef detail::edge_desc_impl<Directed,Vertex,Property> key_type;
    typedef thrust::lvalue_property_map_tag category;
		
    /*!
     *
     */
    __host__ __device__
    adj_list_edge_property_map() { }
    
    /*!
     *
     */
    __host__
    inline Reference
    operator[]( key_type& e) const;
    
    /*!
     *
     */
    __host__
    inline Reference
    operator()( key_type& e) const;
		
		
    /*!
     *
     */
    __host__
    inline void
    put( key_type e, const value_type& v, 
	 thrust::called_from_host_tag);

    /*!
     *
     */
    __device__
    inline void
    put( key_type e, const value_type& v, thrust::called_from_device_tag);
		
    /*template <typename CalledSpace>
			__host__ __device__
			inline void
			put( key_type e, const value_type& v, CalledSpace tag) {
      dispatch_put( e, v, tag);
      }*/
  };
	
	
  namespace detail {
		
    struct adj_list_any_edge_pmap {
      template <typename Graph, typename Property, typename Tag>
      struct bind_ {
				typedef typename property_value<Property,Tag>::type value_type;
				typedef value_type& reference;
				typedef const value_type& const_reference;
				
				typedef adj_list_edge_property_map<typename Graph::directed_category
                                           , value_type
                                           , reference
                                           , typename Graph::vertex_descriptor
                                           , Property
                                           , Tag>  type;
				typedef adj_list_edge_property_map<typename Graph::directed_category
                                           , value_type
                                           , const_reference
                                           , typename Graph::vertex_descriptor
                                           , const Property
                                           , Tag> const_type;
      };
    };
		
    template <typename Tag>
    struct adj_list_choose_edge_pmap_helper {
      typedef adj_list_any_edge_pmap type;
    };
    
    template <typename Tag, typename Graph, typename Property>
    struct adj_list_choose_edge_pmap {
      typedef typename adj_list_choose_edge_pmap_helper<Tag>::type Helper;
      typedef typename Helper::template bind_<Graph,Property,Tag> Bind;
      typedef typename Bind::type       type;
      typedef typename Bind::const_type const_type;
    };
		
    struct adj_list_edge_property_selector {
      template <typename Graph, typename Property, typename Tag>
      struct bind_ {
				typedef adj_list_choose_edge_pmap<Tag,Graph,Property> Choice;
				typedef typename Choice::type       type;
				typedef typename Choice::const_type const_type;
      };
    };
  } // namespace detail
  template <>
  struct edge_property_selector<adj_list_tag> {
    typedef detail::adj_list_edge_property_selector type;
  };
}
#include <thrust/graph/detail/adjacency_list.inl>
